Skip to content

Allow public tuning of the cub::DeviceRunLengthEncode::NonTrivialRuns - #9347

Merged
bernhardmgruber merged 1 commit into
NVIDIA:mainfrom
bernhardmgruber:rle_nontriv_public
Jun 12, 2026
Merged

bernhardmgruber merged 1 commit into
NVIDIA:mainfrom
bernhardmgruber:rle_nontriv_public

Conversation

@bernhardmgruber

Copy link
Copy Markdown
Contributor

Fixes: #8577

New public entities

Entity Enumerators / Data members
cub::RleNonTrivialRunsPolicy int threads_per_block
int items_per_thread
BlockLoadAlgorithm load_algorithm
CacheLoadModifier load_modifier
bool store_with_time_slicing
BlockScanAlgorithm scan_algorithm
LookbackDelayPolicy lookback_delay

@coderabbitai

coderabbitai Bot commented Jun 9, 2026 •

Copy link
Copy Markdown
Contributor

Review Change Stack

Note: CodeRabbit is enabled on this repository as a convenience for maintainers
and contributors. Use your best judgment when considering its review comments and
suggestions — a suggested change may be inadequate, unnecessary, or safe to ignore.
Contributors are not expected to address every comment. Human reviews are what
ultimately matter for merging.

Summary

This PR productizes the tuning API for DeviceRunLengthEncode::NonTrivialRuns by introducing a public cub::RleNonTrivialRunsPolicy and making related dispatch/tuning machinery use that public type (fixes #8577). The new policy exposes tunable parameters:

  • threads_per_block, items_per_thread
  • load_algorithm, load_modifier
  • store_with_time_slicing
  • scan_algorithm
  • lookback_delay (LookbackDelayPolicy)

Changes

Public API Additions

  • Added struct cub::RleNonTrivialRunsPolicy (equality/inequality and stream formatter) in cub/cub/device/dispatch/tuning/tuning_rle_non_trivial_runs.cuh.
  • Benchmark selector and tuning functors now return cub::RleNonTrivialRunsPolicy (benchmarks and tests updated).

Deprecations

  • Deprecated aliases added:
    • using rle_non_trivial_runs_policy = RleNonTrivialRunsPolicy (deprecated in-place of the removed internal struct).
    • using AgentRlePolicy = detail::agent_rle_policy<...> is now a CCCL_DEPRECATED alias (agent_rle.cuh) — the concrete AgentRlePolicy struct was refactored into detail::agent_rle_policy.
  • DeviceRleDispatch was annotated deprecated with a message directing users to DeviceRunLengthEncode.

Implementation & Dispatch Updates

  • Dispatch and kernel code updated to use RleNonTrivialRunsPolicy (dispatch_rle.cuh, tuning files).
  • The policy selection logic across compute-capability branches now constructs and returns RleNonTrivialRunsPolicy.
  • Internal naming shifted from delay_constructor.* to lookback_delay.* where appropriate; dispatch constructs delay_constructor_t from policy.lookback_delay fields.

Documentation

  • Added @par Tuning subsection to cub/cub/device/device_run_length_encode.cuh explaining how to tune NonTrivialRuns via a policy selector with examples referencing tests.

Tests

  • Tests updated to use the public policy type:
    • catch2_test_device_run_length_encode_env.cu: added compile-time C2H test verifying RleNonTrivialRunsPolicy is semiregular and an aggregate; constexpr equality checks added.
    • catch2_test_device_run_length_encode_env_api.cu: added env-based API test demonstrating cuda::execution::tune with a RleNonTrivialRunsPolicy selector.
    • catch2_test_device_run_length_encode_non_trivial_runs.cu and related test files updated to return RleNonTrivialRunsPolicy from selectors.
    • catch2_test_device_rle_custom_policy_hub.cu disables deprecation warnings to continue testing deprecated AgentRlePolicy/DeviceRleDispatch usage.

Other

  • Minor documentation/comment cleanup and small refactors to ensure internal delay-constructor types map cleanly to the new lookback_delay policy field.
  • Files touched include dispatch, tuning, agent, device, benchmarks, and multiple tests.

important:

Walkthrough

This PR makes the NonTrivialRuns tuning API public by adding cub::RleNonTrivialRunsPolicy, refactors AgentRlePolicy into detail::agent_rle_policy with a deprecated alias, updates all selector and dispatch code to use the new public type, and adds compile-time and runtime tests and tuning examples.

Changes

RLE NonTrivialRuns Tuning API Productization

Layer / File(s) Summary
Public policy type and tuning documentation
cub/cub/device/dispatch/tuning/tuning_rle_non_trivial_runs.cuh, cub/cub/device/device_run_length_encode.cuh
RleNonTrivialRunsPolicy struct introduced with tuning fields (including lookback_delay), equality ops, stream formatter; deprecated alias preserves legacy name; device header documents tuning examples.
Internal agent policy refactoring
cub/cub/agent/agent_rle.cuh, cub/cub/device/dispatch/tuning/tuning_rle_non_trivial_runs.cuh
Agent policy implementation moved to detail::agent_rle_policy<...> with public AgentRlePolicy deprecated alias; policy_hub updated to use agent_rle_policy.
Policy selector implementation updates
cub/cub/device/dispatch/tuning/tuning_rle_non_trivial_runs.cuh
policy_selector and policy_selector_from_types now construct and return RleNonTrivialRunsPolicy across SM10x/SM100+, SM90, SM80 branches; make_default_policy updated accordingly.
Dispatch and kernel integration
cub/cub/device/dispatch/dispatch_rle.cuh
DeviceRleSweepKernel and dispatch wiring use RleNonTrivialRunsPolicy; delay fields accessed via policy.lookback_delay.*; DeviceRleDispatch annotated deprecated; detail::rle::dispatch active_policy type updated; legacy TODO removed.
Public API usage in benchmarks and policy selectors
cub/benchmarks/bench/run_length_encode/non_trivial_runs.cu, cub/test/catch2_test_device_run_length_encode_non_trivial_runs.cu
Benchmark and test policy selectors updated to return cub::RleNonTrivialRunsPolicy at public boundaries.
Extended test coverage and compile-time validation
cub/test/catch2_test_device_run_length_encode_env.cu, cub/test/catch2_test_device_run_length_encode_env_api.cu, cub/test/catch2_test_device_rle_custom_policy_hub.cu
Adds constexpr tests validating aggregate/semiregular traits for RleNonTrivialRunsPolicy, adds tuning-enabled env API test (cuda::execution::tune), and suppresses deprecation warnings in legacy test.

Assessment against linked issues

Objective Addressed Explanation
Productize the cub::DeviceRunLengthEncode::NonTrivialRuns tuning API (#8577) ✅
Expose a public tuning policy type and document usage (#8577) ✅
Update internal agent/dispatch code and deprecate legacy API (#8577) ✅
Add compile-time and runtime tests for the new API (#8577) ✅

Suggested reviewers

  • shwina
  • gevtushenko
  • elstehle

Comment @coderabbitai help to get the list of available commands and usage tips.

@github-actions

This comment has been minimized.

@coderabbitai coderabbitai Bot left a comment

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Actionable comments posted: 1

🧹 Nitpick comments (4)
cub/cub/device/dispatch/tuning/tuning_rle_non_trivial_runs.cuh (4)

34-34: 💤 Low value

suggestion: Add explicit @brief tag for consistency.

The Doxygen comment lacks an explicit @brief tag. While guidelines specifically require this for functions, adding it for the struct improves consistency.

-//! The tuning policy for `@ref` DeviceRunLengthEncode::NonTrivialRuns.
+//! `@brief` The tuning policy for `@ref` DeviceRunLengthEncode::NonTrivialRuns.

Source: Coding guidelines


392-406: 💤 Low value

suggestion: Consider marking make_default_policy as noexcept.

The function performs only compile-time arithmetic and aggregate initialization. If default_reduce_by_key_delay_constructor_policy is also noexcept (or constexpr-evaluated), this function doesn't throw and should be marked noexcept per guidelines.

Source: Coding guidelines


408-606: 💤 Low value

suggestion: Consider marking operator() as noexcept.

The function branches on compute capability and returns aggregate-initialized policy objects. If all paths are no-throw (aggregate initialization of POD-like fields), mark it noexcept per guidelines.

Source: Coding guidelines


615-627: 💤 Low value

suggestion: Consider marking operator() as noexcept.

Delegates to selector(cc) which performs compile-time selection. If the selector and construction are no-throw, mark noexcept per guidelines.

Source: Coding guidelines


ℹ️ Review info
⚙️ Run configuration

Configuration used: Path: .coderabbit.yaml

Review profile: CHILL

Plan: Enterprise

Run ID: 76847392-186d-42b7-9e0a-b9e9eaca23aa

📥 Commits

Reviewing files that changed from the base of the PR and between 2c0c1a6 and 21026ad.

📒 Files selected for processing (9)
  • cub/benchmarks/bench/run_length_encode/non_trivial_runs.cu
  • cub/cub/agent/agent_rle.cuh
  • cub/cub/device/device_run_length_encode.cuh
  • cub/cub/device/dispatch/dispatch_rle.cuh
  • cub/cub/device/dispatch/tuning/tuning_rle_non_trivial_runs.cuh
  • cub/test/catch2_test_device_rle_custom_policy_hub.cu
  • cub/test/catch2_test_device_run_length_encode_env.cu
  • cub/test/catch2_test_device_run_length_encode_env_api.cu
  • cub/test/catch2_test_device_run_length_encode_non_trivial_runs.cu
✅ Files skipped from review due to trivial changes (2)
  • cub/cub/device/device_run_length_encode.cuh
  • cub/test/catch2_test_device_rle_custom_policy_hub.cu
🚧 Files skipped from review as they are similar to previous changes (6)
  • cub/benchmarks/bench/run_length_encode/non_trivial_runs.cu
  • cub/test/catch2_test_device_run_length_encode_non_trivial_runs.cu
  • cub/test/catch2_test_device_run_length_encode_env.cu
  • cub/cub/agent/agent_rle.cuh
  • cub/test/catch2_test_device_run_length_encode_env_api.cu
  • cub/cub/device/dispatch/dispatch_rle.cuh

Comment on lines +34 to +72
//! The tuning policy for @ref DeviceRunLengthEncode::NonTrivialRuns.
struct RleNonTrivialRunsPolicy
{
int threads_per_block; //!< Number of threads in a CUDA block
int items_per_thread; //!< Number of items processed per thread
BlockLoadAlgorithm load_algorithm; //!< The @ref BlockLoadAlgorithm used for loading items from global memory
CacheLoadModifier load_modifier; //!< The @ref CacheLoadModifier used for loading items from global memory
bool store_with_time_slicing; //!< Whether to time-slice shared memory for store transpositions
BlockScanAlgorithm scan_algorithm; //!< The @ref BlockScanAlgorithm used for scanning

//! The @ref LookbackDelayPolicy controlling the delay between lookback iterations
LookbackDelayPolicy lookback_delay = {LookbackDelayAlgorithm::fixed_delay, 350, 450};

[[nodiscard]] _CCCL_HOST_DEVICE_API constexpr friend bool
operator==(const RleNonTrivialRunsPolicy& lhs, const RleNonTrivialRunsPolicy& rhs) noexcept
{
return lhs.threads_per_block == rhs.threads_per_block && lhs.items_per_thread == rhs.items_per_thread
&& lhs.load_algorithm == rhs.load_algorithm && lhs.load_modifier == rhs.load_modifier
&& lhs.store_with_time_slicing == rhs.store_with_time_slicing && lhs.scan_algorithm == rhs.scan_algorithm
&& lhs.lookback_delay == rhs.lookback_delay;
}

[[nodiscard]] _CCCL_HOST_DEVICE_API constexpr friend bool
operator!=(const RleNonTrivialRunsPolicy& lhs, const RleNonTrivialRunsPolicy& rhs) noexcept
{
return !(lhs == rhs);
}

#if _CCCL_HOSTED()
friend ::std::ostream& operator<<(::std::ostream& os, const RleNonTrivialRunsPolicy& p)
{
return os
<< "RleNonTrivialRunsPolicy { .threads_per_block = " << p.threads_per_block
<< ", .items_per_thread = " << p.items_per_thread << ", .load_algorithm = " << p.load_algorithm
<< ", .load_modifier = " << p.load_modifier << ", .store_with_time_slicing = " << p.store_with_time_slicing
<< ", .scan_algorithm = " << p.scan_algorithm << ", .lookback_delay = " << p.lookback_delay << " }";
}
#endif // _CCCL_HOSTED()
};

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

⚠️ Potential issue | 🟠 Major

🧩 Analysis chain

🏁 Script executed:

#!/bin/bash
# Search for other public policy struct naming patterns in CUB device headers
rg -n "^struct [A-Z][a-zA-Z]*Policy" cub/cub/device/ --type=cpp -C2

Repository: NVIDIA/cccl

Length of output: 37


Fix CCCL naming/Doxygen compliance in cub/cub/device/dispatch/tuning/tuning_rle_non_trivial_runs.cuh

  • RleNonTrivialRunsPolicy member names (threads_per_block, items_per_thread, store_with_time_slicing, lookback_delay) don’t use the required trailing _ postfix for data members.
  • The Doxygen comment for RleNonTrivialRunsPolicy lacks //! @brief`` (it only says “The tuning policy…”).
  • Add noexcept to policy_selector::make_default_policy(...) and policy_selector::operator()( ::cuda::compute_capability ) const when they don’t throw.

Source: Coding guidelines

@github-actions

Copy link
Copy Markdown
Contributor

🥳 CI Workflow Results

🟩 Finished in 2h 43m: Pass: 100%/287 | Total: 12d 11h | Max: 2h 42m | Hits: 15%/1399288

See results here.

@bernhardmgruber
bernhardmgruber merged commit d307271 into NVIDIA:main Jun 12, 2026
309 checks passed
@github-project-automation github-project-automation Bot moved this from In Review to Done in CCCL Jun 12, 2026
@bernhardmgruber
bernhardmgruber deleted the rle_nontriv_public branch June 12, 2026 21:45
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

None yet

Projects

Archived in project

Development

Successfully merging this pull request may close these issues.

Productize the cub::DeviceRunLengthEncode::NonTrivialRuns tuning API

2 participants